#include <cuda_runtime.h>
#include <transpose_ops.h>
template<typename T>
__global__ void transpose_kernel(T* d_out,T* d_in,const unsigned int rows,const unsigned int cols){
    auto idx = blockIdx.x*blockDim.x+threadIdx.x;
    auto idy = blockIdx.y*blockDim.y+threadIdx.y;
    d_out[idx*rows+idy] = d_in[idy*cols+idx];

}
template <typename T>
void transpose_cuda(T *h_out, T *h_in, const unsigned int rows,
                    const unsigned int cols){
    const int N = rows*cols;
    size_t size = sizeof(T)*N;
    T*d_in,*d_out;
    cudaMalloc(reinterpret_cast<void**>(&d_in),size);
    cudaMalloc(reinterpret_cast<void**>(&d_out),size);
    cudaMemcpy(d_in, h_in, size, cudaMemcpyHostToDevice);
    dim3 block = {32,32};
    dim3 grid = {(rows+block.x-1)/block.x,(cols+block.y-1)/block.y};
    transpose_kernel<<<grid,block>>>(d_out, d_in, rows, cols);
    cudaMemcpy(h_out,d_out,size,cudaMemcpyDeviceToHost);
    cudaFree(d_out);
    cudaFree(d_in);
    
}
template void transpose_cuda<int>(int*, int *, const unsigned int, const unsigned int);
template void transpose_cuda<float>(float*, float *, const unsigned int, const unsigned int);
